#include<stdio.h>
#include<stdlib.h>
#include"fftz.h"

static const int N = 16777216;
static const int stride = 2097152;

int main(int argc, char** argv){
	CLEnv cl;
	CLKernel ke1;
	CLKernel ke2;
	cl_mem buffer_data0 = NULL;
	cl_mem buffer_data1 = NULL;
	size_t globalws[1];
	size_t localws[1];
	cl_event myevent;
	cl_ulong task_queued, task_start, task_end;
	double time1, time2, speed1, speed2, size;
	int err;
	
	if(initCLEnv(&cl) < 0){
		printf("Initialization of CLEvn failed, program exiting..\n");
		return -1;
	}
	
	if(initCLKernel(&cl, &ke1, "stream.cl", "writeOnce") < 0){
		printf("Initialization of ke1 failed, program exiting..\n");
		releaseCLEnv(&cl);
		return -1;
	}
	
	if(initCLKernel(&cl, &ke2, "stream.cl", "writeTwice") < 0){
		printf("Initialization of ke2 failed, program exiting..\n");
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return -1;
	}
	
	buffer_data0 = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, sizeof(float) * N, NULL, &err);
	if(err < 0){
		printf("Runtime Error: gpuFFTBase. Calling clCreateBuffer failed with code %d.\n", err);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}
	buffer_data1 = clCreateBuffer(cl.context, CL_MEM_READ_WRITE, sizeof(float) * N * 2, NULL, &err);
	if(err < 0){
		printf("Runtime Error: gpuFFTBase. Calling clCreateBuffer failed with code %d.\n", err);
		clReleaseMemObject(buffer_data0);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}
	
	//run the first kernel
	err += clSetKernelArg(ke1.kernel, 0, sizeof(cl_mem), &buffer_data0);
	err += clSetKernelArg(ke1.kernel, 1, sizeof(cl_mem), &buffer_data1);
	err += clSetKernelArg(ke1.kernel, 2, sizeof(int), &stride);
	err += clSetKernelArg(ke1.kernel, 3, sizeof(int), &N);
	if(err < 0){
		printf("Runtime Error: Setting args of ke1 failed with code %d.\n", err);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}

	globalws[0] = N / 8;
	localws[0] = ke1.groupSize;
	err = clEnqueueNDRangeKernel(cl.queue, ke1.kernel, 1, NULL, globalws, localws, 0, NULL, &myevent);
	if(err < 0){
		printf("Runtime Error: Enqueuing ke1 failed with code %d.\n", err);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}
	clFinish(cl.queue);
	clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &task_start, NULL);
	clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &task_end, NULL);
	time1 = (double)(task_end - task_start) / 1000.0;	//us
	clReleaseEvent(myevent);
	
	//run the second kernel
	err += clSetKernelArg(ke2.kernel, 0, sizeof(cl_mem), &buffer_data0);
	err += clSetKernelArg(ke2.kernel, 1, sizeof(cl_mem), &buffer_data1);
	err += clSetKernelArg(ke2.kernel, 2, sizeof(int), &stride);
	err += clSetKernelArg(ke2.kernel, 3, sizeof(int), &N);
	if(err < 0){
		printf("Runtime Error: Setting args of ke2 failed with code %d.\n", err);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}

	globalws[0] = N / 8;
	localws[0] = ke2.groupSize;
	err = clEnqueueNDRangeKernel(cl.queue, ke2.kernel, 1, NULL, globalws, localws, 0, NULL, &myevent);
	if(err < 0){
		printf("Runtime Error: Enqueuing ke2 failed with code %d.\n", err);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		releaseCLKernel(&ke2);
		releaseCLKernel(&ke1);
		releaseCLEnv(&cl);
		return err;
	}
	clFinish(cl.queue);
	clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &task_start, NULL);
	clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &task_end, NULL);
	time2 = (double)(task_end - task_start) / 1000.0;	//us
	clReleaseEvent(myevent);
	
	printf("Result:\n");
	printf("  time1: %.0f us\n", time1);
	printf("  time2: %.0f us\n", time2);
	size = (double)(N * 4 / 1024 / 1024);
	speed2 = size / (time2 - time1) * 1000000;
	speed1 = size / (time1 / 1000000 - size / speed2);
	printf("  speed1: %.0f MB/s\n", speed1);
	printf("  speed2: %.0f MB/s\n", speed2);
	
	clReleaseMemObject(buffer_data1);
	clReleaseMemObject(buffer_data0);
	releaseCLKernel(&ke2);
	releaseCLKernel(&ke1);
	releaseCLEnv(&cl);
}